#include "kernel_operator.h"
#include <math.h>
using namespace AscendC;
constexpr int32_t BUFFER_NUM = 1;
class KernelSpence
{
public:
  __aicore__ inline KernelSpence() {}
  __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, uint32_t blockLength,
                              uint32_t tileNum, uint32_t tileLength,
                              uint32_t lasttileLength, uint32_t typeKey)
  {
    ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");
    this->typeKey = typeKey;
    this->blockLength = blockLength;
    this->tileNum =
        tileNum ASSERT(tileNum != 0 && "tile num can not be zero!");
    this->tileLength = tileLength / BUFFER_NUM;
    this->lasttileLength = lasttileLength;
    if(this->tileNum == 1 && BUFFER_NUM == 1){
      uint32_t ALIGN_NUM = 32;
      if(typeKey == 0) ALIGN_NUM = 8;
      else ALIGN_NUM = 16;
      if(lasttileLength % ALIGN_NUM == 0){
        this->tileLength = lasttileLength;
      }else{
        this->tileLength = (lasttileLength + ALIGN_NUM - 1) / ALIGN_NUM * ALIGN_NUM;
      }
    }
    xGm.SetGlobalBuffer((__gm__ DTYPE_X *)x,
                        this->blockLength);
    yGm.SetGlobalBuffer((__gm__ DTYPE_X *)y,
                          this->blockLength);
    
    this->PI2O6 = (float)(M_PI * M_PI / 6);
    this->A[0] = 4.65128586073990045278E-5;
    this->A[1] = 7.31589045238094711071E-3;
    this->A[2] = 1.33847639578309018650E-1;
    this->A[3] = 8.79691311754530315341E-1;
    this->A[4] = 2.71149851196553469920E0;
    this->A[5] = 4.25697156008121755724E0;
    this->A[6] = 3.29771340985225106936E0;
    this->A[7] = 1.00000000000000000126E0;
    this->B[0] = 6.90990488912553276999E-4;
    this->B[1] = 2.54043763932544379113E-2;
    this->B[2] = 2.82974860602568089943E-1;
    this->B[3] = 1.41172597751831069617E0;
    this->B[4] = 3.63800533345137075418E0;
    this->B[5] = 5.03278880143316990390E0;
    this->B[6] = 3.54771340985225096217E0;
    this->B[7] = 9.99999999999999998740E-1;
    this->nanValue = NAN;

    pipe.InitBuffer(inQueueX, BUFFER_NUM, this->tileLength * sizeof(DTYPE_X));
    pipe.InitBuffer(outQueueY, BUFFER_NUM, this->tileLength * sizeof(DTYPE_X));
    pipe.InitBuffer(compareBuf, this->tileLength * sizeof(uint8_t));
    pipe.InitBuffer(calcBuf, this->tileLength * sizeof(float));
    pipe.InitBuffer(calcBufzero, this->tileLength * sizeof(float));
    pipe.InitBuffer(calcBufone, this->tileLength * sizeof(float));
    pipe.InitBuffer(calcBuftwo, this->tileLength * sizeof(float));
    pipe.InitBuffer(calcBufhalf, this->tileLength * sizeof(float));
    pipe.InitBuffer(calcBufthreehalf, this->tileLength * sizeof(float));
    if(typeKey == 1){// half
      pipe.InitBuffer(xfBuf, this->tileLength * sizeof(float));
      pipe.InitBuffer(yfBuf, this->tileLength * sizeof(float));
      masklen = 128;
      if(this->tileLength < masklen){
        masklen = this->tileLength;
      }else{
        while(this->tileLength % masklen != 0){
          masklen /= 2;
        }
        if(masklen < 1)masklen = 1;
      }
    }else{// float  bfloat16->float
      if(typeKey != 0){
        pipe.InitBuffer(xfBuf, BUFFER_NUM * this->tileLength * sizeof(float));
        pipe.InitBuffer(yfBuf, BUFFER_NUM * this->tileLength * sizeof(float));
      }
      masklen = 64;
      if(this->tileLength < masklen){
        masklen = this->tileLength;
      }else{
        while(this->tileLength % masklen != 0){
          masklen /= 2;
        }
        if(masklen < 1)masklen = 1;
      }
    }
  }
  __aicore__ inline void Process()
  {
    int32_t loopCount = this->tileNum * BUFFER_NUM;
    for (int32_t i = 0; i < loopCount; i++)
    {
      CopyIn(i);
      Compute(i);
      CopyOut(i);
    }
  }

private:
  __aicore__ inline void CopyIn(int32_t progress)
  {
    LocalTensor<DTYPE_X> xLocal = inQueueX.AllocTensor<DTYPE_X>();

    if (BUFFER_NUM == 1) {
      if (progress == this->tileNum - 1) {
        if (progress == 0) {
          DataCopy(xLocal[0], xGm[0], this->tileLength);
        } else {
          DataCopy(
              xLocal[0],
              xGm[(progress - 1) * this->tileLength + this->lasttileLength],
              this->tileLength);
        }
      } else {
        DataCopy(xLocal[0], xGm[progress * this->tileLength],
                 this->tileLength);
        
      }
    }
    if (BUFFER_NUM == 2) {
      if ((progress == (this->tileNum * BUFFER_NUM - 2)) ||
          (progress == (this->tileNum * BUFFER_NUM - 1))) {
        DataCopy(
            xLocal[0],
            xGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
            (this->tileLength));

      }
      else {
        DataCopy(xLocal[0], xGm[progress * (this->tileLength)],
                 (this->tileLength));
      }
    }
    inQueueX.EnQue(xLocal);
    
  }
  __aicore__ inline void Compute(int32_t progress)
  {
    BinaryRepeatParams repeatParams = { 1, 1, 1, 8, 8, 8 };
    if(this->typeKey == 0){ // float
      LocalTensor<float> xLocal = inQueueX.DeQue<float>();
      LocalTensor<float> yLocal = outQueueY.AllocTensor<float>();
      LocalTensor<float> zeroLocal = calcBufzero.Get<float>();
      LocalTensor<float> oneLocal = calcBufone.Get<float>();
      LocalTensor<float> halfLocal = calcBufhalf.Get<float>();
      LocalTensor<float> twoLocal = calcBuftwo.Get<float>();
      LocalTensor<float> wLocal = calcBuf.Get<float>();
      LocalTensor<float> threehalfLocal = calcBufthreehalf.Get<float>();
      LocalTensor<uint8_t> resLocal = compareBuf.Get<uint8_t>();
      LocalTensor<float> oneOxLocal = twoLocal;
      LocalTensor<float> neyLocal = twoLocal;
      LocalTensor<float> oneOySyLocal = twoLocal;
      LocalTensor<float> wresALocal = zeroLocal;
      LocalTensor<float> wresBLocal = twoLocal;
      // compute
      
      Duplicate(halfLocal, (float)0.5, this->tileLength);
      Duplicate(twoLocal, (float)2.0, this->tileLength);
      Duplicate(threehalfLocal, (float)1.5, this->tileLength);
      // xLocal 不能动，在最后需要和zeroLocal与oneLocal比较，更新yLocal
      Duplicate(oneLocal, (float)1.0, this->tileLength);
      
      Compare(resLocal, xLocal, twoLocal, CMPMODE::LT, this->tileLength);
      Div(oneOxLocal, oneLocal, xLocal, this->tileLength);
      Select(yLocal, resLocal, xLocal, oneOxLocal, SELMODE::VSEL_TENSOR_TENSOR_MODE, masklen, tileLength / masklen, repeatParams);

      Compare(resLocal, yLocal, halfLocal, CMPMODE::LT, this->tileLength); // y < half
      Muls(neyLocal, yLocal, (float)-1.0, this->tileLength);
      Adds(yLocal, yLocal, (float)-1.0, this->tileLength);
      Select(wLocal, resLocal, neyLocal, yLocal, SELMODE::VSEL_TENSOR_TENSOR_MODE, masklen, tileLength / masklen, repeatParams);
      Adds(yLocal, yLocal, (float)1.0, this->tileLength);

      Compare(resLocal, threehalfLocal, yLocal, CMPMODE::LT, this->tileLength);
      Div(oneOySyLocal, oneLocal, yLocal, this->tileLength);
      Adds(oneOySyLocal, oneOySyLocal, (float)-1.0, this->tileLength);
      Select(wLocal, resLocal, oneOySyLocal, wLocal, SELMODE::VSEL_TENSOR_TENSOR_MODE, masklen, tileLength / masklen, repeatParams);

      // 多项式处理
      Muls(wresALocal, wLocal, (float)A[0], this->tileLength);
      Adds(wresALocal, wresALocal, (float)A[1], this->tileLength);
      Muls(wresBLocal, wLocal, (float)B[0], this->tileLength);
      Adds(wresBLocal, wresBLocal, (float)B[1], this->tileLength);
      for(int i = 2; i < 8; i++){
        Mul(wresALocal, wresALocal, wLocal, this->tileLength);
        Adds(wresALocal, wresALocal, (float)A[i], this->tileLength);
        Mul(wresBLocal, wresBLocal, wLocal, this->tileLength);
        Adds(wresBLocal, wresBLocal, (float)B[i], this->tileLength);
      }

      Div(wresALocal, wresALocal, wresBLocal, this->tileLength);
      Muls(wresBLocal, wLocal, (float)-1.0, this->tileLength);
      Mul(wLocal, wresBLocal, wresALocal, this->tileLength); // spence ---- wLocal

      Ln(wresALocal, yLocal, this->tileLength); // z --- wresALocal --- zeroLocal

      Compare(resLocal, yLocal, halfLocal, CMPMODE::LT, this->tileLength);
      LocalTensor<float> tmpLocal = twoLocal;
      Sub(tmpLocal, oneLocal, yLocal, this->tileLength);
      Ln(tmpLocal, tmpLocal, this->tileLength);
      Mul(tmpLocal, tmpLocal, wresALocal, this->tileLength);
      Muls(tmpLocal, tmpLocal, (float)-1.0, this->tileLength);
      Adds(tmpLocal, tmpLocal, (float)PI2O6, this->tileLength);
      Sub(tmpLocal, tmpLocal, wLocal, this->tileLength);
      Select(wLocal, resLocal, tmpLocal, wLocal, SELMODE::VSEL_TENSOR_TENSOR_MODE, masklen, tileLength / masklen, repeatParams);
      
      Compare(resLocal, threehalfLocal, xLocal, CMPMODE::LT, this->tileLength);
      Mul(tmpLocal, wresALocal, wresALocal, this->tileLength);
      Muls(tmpLocal, tmpLocal, (float)-0.5, this->tileLength);
      Sub(tmpLocal, tmpLocal, wLocal, this->tileLength);
      Select(wLocal, resLocal, tmpLocal, wLocal, SELMODE::VSEL_TENSOR_TENSOR_MODE, masklen, tileLength / masklen, repeatParams);
      

      // 三类特殊值
      Duplicate(zeroLocal, (float)0.0, this->tileLength);
      Compare(resLocal, xLocal, oneLocal, CMPMODE::NE, this->tileLength);
      Select(yLocal, resLocal, wLocal, (float)0.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, this->tileLength);//masklen, tileLength / masklen, repeatParams);

      Compare(resLocal, xLocal, zeroLocal, CMPMODE::GE, this->tileLength);
      Select(yLocal, resLocal, yLocal, (float)NAN, SELMODE::VSEL_TENSOR_SCALAR_MODE, this->tileLength);//masklen, tileLength / masklen, repeatParams);

      Compare(resLocal, xLocal, zeroLocal, CMPMODE::NE, this->tileLength);
      Select(yLocal, resLocal, yLocal, (float)this->PI2O6, SELMODE::VSEL_TENSOR_SCALAR_MODE, this->tileLength);//masklen, tileLength / masklen, repeatParams);

      outQueueY.EnQue<float>(yLocal);
      inQueueX.FreeTensor(xLocal);
    }else if(this->typeKey == 1){ // half
      LocalTensor<half> xLocal_h = inQueueX.DeQue<half>();
      LocalTensor<half> yLocal_h = outQueueY.AllocTensor<half>();
      LocalTensor<float> xLocal = xfBuf.Get<float>();
      LocalTensor<float> yLocal = yfBuf.Get<float>();
      LocalTensor<float> zeroLocal = calcBufzero.Get<float>();
      LocalTensor<float> oneLocal = calcBufone.Get<float>();
      LocalTensor<float> halfLocal = calcBufhalf.Get<float>();
      LocalTensor<float> twoLocal = calcBuftwo.Get<float>();
      LocalTensor<float> wLocal = calcBuf.Get<float>();
      LocalTensor<float> threehalfLocal = calcBufthreehalf.Get<float>();
      LocalTensor<uint8_t> resLocal = compareBuf.Get<uint8_t>();
      LocalTensor<float> oneOxLocal = twoLocal;
      LocalTensor<float> neyLocal = twoLocal;
      LocalTensor<float> oneOySyLocal = twoLocal;
      LocalTensor<float> wresALocal = zeroLocal;
      LocalTensor<float> wresBLocal = twoLocal;
      // compute
      Cast(xLocal, xLocal_h, RoundMode::CAST_NONE, this->tileLength);
      Duplicate(halfLocal, (float)0.5, this->tileLength);
      Duplicate(twoLocal, (float)2.0, this->tileLength);
      Duplicate(threehalfLocal, (float)1.5, this->tileLength);
      // xLocal 不能动，在最后需要和zeroLocal与oneLocal比较，更新yLocal
      Duplicate(oneLocal, (float)1.0, this->tileLength);
      
      Compare(resLocal, xLocal, twoLocal, CMPMODE::LT, this->tileLength);
      Div(oneOxLocal, oneLocal, xLocal, this->tileLength);
      Select(yLocal, resLocal, xLocal, oneOxLocal, SELMODE::VSEL_TENSOR_TENSOR_MODE, masklen, tileLength / masklen, repeatParams);

      Compare(resLocal, yLocal, halfLocal, CMPMODE::LT, this->tileLength); // y < half
      Muls(neyLocal, yLocal, (float)-1.0, this->tileLength);
      Adds(yLocal, yLocal, (float)-1.0, this->tileLength);
      Select(wLocal, resLocal, neyLocal, yLocal, SELMODE::VSEL_TENSOR_TENSOR_MODE, masklen, tileLength / masklen, repeatParams);
      Adds(yLocal, yLocal, (float)1.0, this->tileLength);

      Compare(resLocal, threehalfLocal, yLocal, CMPMODE::LT, this->tileLength);
      Div(oneOySyLocal, oneLocal, yLocal, this->tileLength);
      Adds(oneOySyLocal, oneOySyLocal, (float)-1.0, this->tileLength);
      Select(wLocal, resLocal, oneOySyLocal, wLocal, SELMODE::VSEL_TENSOR_TENSOR_MODE, masklen, tileLength / masklen, repeatParams);

      // 多项式处理
      Muls(wresALocal, wLocal, (float)A[0], this->tileLength);
      Adds(wresALocal, wresALocal, (float)A[1], this->tileLength);
      Muls(wresBLocal, wLocal, (float)B[0], this->tileLength);
      Adds(wresBLocal, wresBLocal, (float)B[1], this->tileLength);
      for(int i = 2; i < 8; i++){
        Mul(wresALocal, wresALocal, wLocal, this->tileLength);
        Adds(wresALocal, wresALocal, (float)A[i], this->tileLength);
        Mul(wresBLocal, wresBLocal, wLocal, this->tileLength);
        Adds(wresBLocal, wresBLocal, (float)B[i], this->tileLength);
      }

      Div(wresALocal, wresALocal, wresBLocal, this->tileLength);
      Muls(wresBLocal, wLocal, (float)-1.0, this->tileLength);
      Mul(wLocal, wresBLocal, wresALocal, this->tileLength); // spence ---- wLocal

      Ln(wresALocal, yLocal, this->tileLength); // z --- wresALocal --- zeroLocal

      Compare(resLocal, yLocal, halfLocal, CMPMODE::LT, this->tileLength);
      LocalTensor<float> tmpLocal = twoLocal;
      Sub(tmpLocal, oneLocal, yLocal, this->tileLength);
      Ln(tmpLocal, tmpLocal, this->tileLength);
      Mul(tmpLocal, tmpLocal, wresALocal, this->tileLength);
      Muls(tmpLocal, tmpLocal, (float)-1.0, this->tileLength);
      Adds(tmpLocal, tmpLocal, (float)PI2O6, this->tileLength);
      Sub(tmpLocal, tmpLocal, wLocal, this->tileLength);
      Select(wLocal, resLocal, tmpLocal, wLocal, SELMODE::VSEL_TENSOR_TENSOR_MODE, masklen, tileLength / masklen, repeatParams);
      
      Compare(resLocal, threehalfLocal, xLocal, CMPMODE::LT, this->tileLength);
      Mul(tmpLocal, wresALocal, wresALocal, this->tileLength);
      Muls(tmpLocal, tmpLocal, (float)-0.5, this->tileLength);
      Sub(tmpLocal, tmpLocal, wLocal, this->tileLength);
      Select(wLocal, resLocal, tmpLocal, wLocal, SELMODE::VSEL_TENSOR_TENSOR_MODE, masklen, tileLength / masklen, repeatParams);
      

      // 三类特殊值
      Duplicate(zeroLocal, (float)0.0, this->tileLength);
      Compare(resLocal, xLocal, oneLocal, CMPMODE::NE, this->tileLength);
      Select(yLocal, resLocal, wLocal, (float)0.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, this->tileLength);//masklen, tileLength / masklen, repeatParams);

      Compare(resLocal, xLocal, zeroLocal, CMPMODE::GE, this->tileLength);
      Select(yLocal, resLocal, yLocal, (float)NAN, SELMODE::VSEL_TENSOR_SCALAR_MODE, this->tileLength);//masklen, tileLength / masklen, repeatParams);

      Compare(resLocal, xLocal, zeroLocal, CMPMODE::NE, this->tileLength);
      Select(yLocal, resLocal, yLocal, (float)this->PI2O6, SELMODE::VSEL_TENSOR_SCALAR_MODE, this->tileLength);//masklen, tileLength / masklen, repeatParams);

      Cast(yLocal_h, yLocal, RoundMode::CAST_NONE, this->tileLength);
      outQueueY.EnQue<half>(yLocal_h);
      inQueueX.FreeTensor(xLocal_h);
    }else{ //bfloat16_t cast to float
      LocalTensor<bfloat16_t> xLocal_h = inQueueX.DeQue<bfloat16_t>();
      LocalTensor<bfloat16_t> yLocal_h = outQueueY.AllocTensor<bfloat16_t>();
      LocalTensor<float> xLocal = xfBuf.Get<float>();
      LocalTensor<float> yLocal = yfBuf.Get<float>();
      LocalTensor<float> zeroLocal = calcBufzero.Get<float>();
      LocalTensor<float> oneLocal = calcBufone.Get<float>();
      LocalTensor<float> halfLocal = calcBufhalf.Get<float>();
      LocalTensor<float> twoLocal = calcBuftwo.Get<float>();
      LocalTensor<float> wLocal = calcBuf.Get<float>();
      LocalTensor<float> threehalfLocal = calcBufthreehalf.Get<float>();
      LocalTensor<uint8_t> resLocal = compareBuf.Get<uint8_t>();
      LocalTensor<float> oneOxLocal = twoLocal;
      LocalTensor<float> neyLocal = twoLocal;
      LocalTensor<float> oneOySyLocal = twoLocal;
      LocalTensor<float> wresALocal = zeroLocal;
      LocalTensor<float> wresBLocal = twoLocal;
      // compute
      Cast(xLocal, xLocal_h, RoundMode::CAST_NONE, this->tileLength);
      Duplicate(halfLocal, (float)0.5, this->tileLength);
      Duplicate(twoLocal, (float)2.0, this->tileLength);
      Duplicate(threehalfLocal, (float)1.5, this->tileLength);
      // xLocal 不能动，在最后需要和zeroLocal与oneLocal比较，更新yLocal
      Duplicate(oneLocal, (float)1.0, this->tileLength);
      
      Compare(resLocal, xLocal, twoLocal, CMPMODE::LT, this->tileLength);
      Div(oneOxLocal, oneLocal, xLocal, this->tileLength);
      Select(yLocal, resLocal, xLocal, oneOxLocal, SELMODE::VSEL_TENSOR_TENSOR_MODE, masklen, tileLength / masklen, repeatParams);

      Compare(resLocal, yLocal, halfLocal, CMPMODE::LT, this->tileLength); // y < half
      Muls(neyLocal, yLocal, (float)-1.0, this->tileLength);
      Adds(yLocal, yLocal, (float)-1.0, this->tileLength);
      Select(wLocal, resLocal, neyLocal, yLocal, SELMODE::VSEL_TENSOR_TENSOR_MODE, masklen, tileLength / masklen, repeatParams);
      Adds(yLocal, yLocal, (float)1.0, this->tileLength);

      Compare(resLocal, threehalfLocal, yLocal, CMPMODE::LT, this->tileLength);
      Div(oneOySyLocal, oneLocal, yLocal, this->tileLength);
      Adds(oneOySyLocal, oneOySyLocal, (float)-1.0, this->tileLength);
      Select(wLocal, resLocal, oneOySyLocal, wLocal, SELMODE::VSEL_TENSOR_TENSOR_MODE, masklen, tileLength / masklen, repeatParams);

      // 多项式处理
      Muls(wresALocal, wLocal, (float)A[0], this->tileLength);
      Adds(wresALocal, wresALocal, (float)A[1], this->tileLength);
      Muls(wresBLocal, wLocal, (float)B[0], this->tileLength);
      Adds(wresBLocal, wresBLocal, (float)B[1], this->tileLength);
      for(int i = 2; i < 8; i++){
        Mul(wresALocal, wresALocal, wLocal, this->tileLength);
        Adds(wresALocal, wresALocal, (float)A[i], this->tileLength);
        Mul(wresBLocal, wresBLocal, wLocal, this->tileLength);
        Adds(wresBLocal, wresBLocal, (float)B[i], this->tileLength);
      }

      Div(wresALocal, wresALocal, wresBLocal, this->tileLength);
      Muls(wresBLocal, wLocal, (float)-1.0, this->tileLength);
      Mul(wLocal, wresBLocal, wresALocal, this->tileLength); // spence ---- wLocal

      Ln(wresALocal, yLocal, this->tileLength); // z --- wresALocal --- zeroLocal

      Compare(resLocal, yLocal, halfLocal, CMPMODE::LT, this->tileLength);
      LocalTensor<float> tmpLocal = twoLocal;
      Sub(tmpLocal, oneLocal, yLocal, this->tileLength);
      Ln(tmpLocal, tmpLocal, this->tileLength);
      Mul(tmpLocal, tmpLocal, wresALocal, this->tileLength);
      Muls(tmpLocal, tmpLocal, (float)-1.0, this->tileLength);
      Adds(tmpLocal, tmpLocal, (float)PI2O6, this->tileLength);
      Sub(tmpLocal, tmpLocal, wLocal, this->tileLength);
      Select(wLocal, resLocal, tmpLocal, wLocal, SELMODE::VSEL_TENSOR_TENSOR_MODE, masklen, tileLength / masklen, repeatParams);
      
      Compare(resLocal, threehalfLocal, xLocal, CMPMODE::LT, this->tileLength);
      Mul(tmpLocal, wresALocal, wresALocal, this->tileLength);
      Muls(tmpLocal, tmpLocal, (float)-0.5, this->tileLength);
      Sub(tmpLocal, tmpLocal, wLocal, this->tileLength);
      Select(wLocal, resLocal, tmpLocal, wLocal, SELMODE::VSEL_TENSOR_TENSOR_MODE, masklen, tileLength / masklen, repeatParams);
      

      // 三类特殊值
      Duplicate(zeroLocal, (float)0.0, this->tileLength);
      Compare(resLocal, xLocal, oneLocal, CMPMODE::NE, this->tileLength);
      Select(yLocal, resLocal, wLocal, (float)0.0, SELMODE::VSEL_TENSOR_SCALAR_MODE, this->tileLength);//masklen, tileLength / masklen, repeatParams);

      Compare(resLocal, xLocal, zeroLocal, CMPMODE::GE, this->tileLength);
      Select(yLocal, resLocal, yLocal, (float)NAN, SELMODE::VSEL_TENSOR_SCALAR_MODE, this->tileLength);//masklen, tileLength / masklen, repeatParams);

      Compare(resLocal, xLocal, zeroLocal, CMPMODE::NE, this->tileLength);
      Select(yLocal, resLocal, yLocal, (float)this->PI2O6, SELMODE::VSEL_TENSOR_SCALAR_MODE, this->tileLength);//masklen, tileLength / masklen, repeatParams);

      Cast(yLocal_h, yLocal, RoundMode::CAST_NONE, this->tileLength);
      outQueueY.EnQue<bfloat16_t>(yLocal_h);
      inQueueX.FreeTensor(xLocal_h);
    }
  }
  __aicore__ inline void CopyOut(int32_t progress)
  {
    LocalTensor<DTYPE_X> yLocal = outQueueY.DeQue<DTYPE_X>();

    if (BUFFER_NUM == 1) {
      if (progress == this->tileNum - 1) {
        if (progress == 0) {
          DataCopy(yGm[0], yLocal, this->tileLength);
        } else {
          DataCopy(
              yGm[(progress - 1) * this->tileLength + this->lasttileLength],
              yLocal, this->tileLength);
        }
      } else {
        DataCopy(yGm[progress * this->tileLength], yLocal,
                 this->tileLength);
      }
    }
    if (BUFFER_NUM == 2) {
      if ((progress == (this->tileNum * BUFFER_NUM - 2)) ||
          (progress == (this->tileNum * BUFFER_NUM - 1))) {
        DataCopy(
            yGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
            yLocal, (this->tileLength));
      }
      else {
        DataCopy(yGm[progress * (this->tileLength)], yLocal,
                 (this->tileLength));
      }
    }

    outQueueY.FreeTensor(yLocal);
  }

private:
  TPipe pipe;
  TQue<QuePosition::VECIN, BUFFER_NUM> inQueueX;
  TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY;
  TBuf<QuePosition::VECCALC> calcBuf, compareBuf, calcBuftwo, calcBufone, calcBufzero, calcBufthreehalf, calcBufhalf, xfBuf, yfBuf;
  GlobalTensor<DTYPE_X> xGm;
  GlobalTensor<DTYPE_Y> yGm;
  uint32_t blockLength;
  uint32_t tileNum;
  uint32_t tileLength;
  uint32_t lasttileLength;
  uint32_t typeKey;
  uint32_t masklen;
  float A[8], B[8], PI2O6, nanValue;
};
extern "C" __global__ __aicore__ void spence(GM_ADDR x, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling) {
    GET_TILING_DATA(tiling_data, tiling);
    KernelSpence op;


    op.Init(x, y, tiling_data.blockLength,
            tiling_data.tileNum, tiling_data.tileLength,
            tiling_data.lasttileLength, tiling_data.typeKey);
    op.Process();
    // TODO: user kernel impl
}
#ifndef __CCE_KT_TEST__
void spence_do(uint32_t blockDim, void* l2ctrl, void* stream,
                       uint8_t* x, uint8_t* y,
                       uint8_t* workspace, uint8_t* tiling) {
  spence<<<blockDim, l2ctrl, stream>>>(x, y, workspace, tiling);
}
#endif